module test_cuda_03_kernels

  implicit none

  integer, parameter :: n = 100
  real(8) unified_array(n)

contains

  attributes(global) subroutine kernel(x)

    real(8) x(n)

    integer i

    i = threadIdx%x

    x(i) = unified_array(i)

  end subroutine kernel

end module test_cuda_03_kernels

program test_cuda_03

  use cudafor
  use test_cuda_03_kernels

  implicit none

  real(8) x(n)
  integer i, ierr

  unified_array = [((i), i = 1, n)]

  call kernel<<<1,n>>>(x)

  ierr = cudaDeviceSynchronize()

  write(*, *) sum(x)

end program test_cuda_03
